#include <stdio.h>

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

int main(void)
{
  const int n = 64;
  int a[n], r[n], d[n];

  for (int i = 0; i < n; i++) {
    a[i] = i;
    r[i] = n-i-1;
    d[i] = 0;
  }

  int *d_d;
  cudaMalloc(&d_d, n * sizeof(int));

  // run version with static shared memory
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  staticReverse<<<1,n>>>(d_d, n);
  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++)
    if (d[i] != r[i]) {
		printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
		exit(-1);
	}
  printf("static success\n");

  // run dynamic shared memory version
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  //your code here
  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
  //end of your code
  cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) {
    if (d[i] != r[i]) {
		printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
		exit(-1);
	}
  printf("dynamic success\n");
    }
}

/*
共享内存的概念
__shared__变量声明说明符
共享内存实际上是可受用户控制的一级缓存。申请共享内存后，其内容在每一个用到的block被复制一遍，使得在每个block内，每一个thread都可以访问和操作这块内存，而无法访问其他block内的共享内存。这种机制就使得一个block之内的所有线程可以互相交流和合作。

在设备代码中声明共享内存要使用__shared__变量声明说明符。

共享内存有两种方法：静态与动态。

如果共享内存数组的大小在编译时就可以确定，我们就可以显式地声明固定大小的数组。

如果共享内存的大小在编译时不能确定，则需要用动态分配共享内存的方式。在这种情况下，每个线程块中共享内存的大小必须在核函数第三个执行配置参数中指定(以字节为单位)。如下所示：

dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);
*/